#ifndef __MEMORY_ROUTINES__
#define __MEMORY_ROUTINES__

//static __device__ void coopExtractRowVector(float* data, int dataPitch, int index, int dimension, float* destination) {
//	float* xiRowPtr = data + (index * dataPitch) + threadIdx.x;
//	for(int currentDim = threadIdx.x; currentDim < dimension; currentDim += blockDim.x) {
//		destination[currentDim] = *xiRowPtr;
//		xiRowPtr += blockDim.x;
//	}
//}

static __device__ void scatterRowVector(unsigned int tid, float* data, int dimension, float* dest) {
	for(int i=tid; i<dimension; i += blockDim.x, data+=blockDim.x) {
		dest[i] = *data;
	}
}

static __device__ void scatterTextureRowVector(int tid, float* dest, int idxX, int dimension) {
	for(int i=tid; i<dimension; i += blockDim.x) {
		dest[i] = 0.0f;
	}
	//__syncthreads();
	__shared__ int len;
	if(tid == 0) {
		len = tex1Dfetch(ellRowLenTex, idxX);
	}
	__syncthreads();
	for(int i=tid; i<len;i+=blockDim.x) {
		int2 X = tex2D(ellDataTex, idxX, i);
		dest[X.x] = __int_as_float(X.y);
	}
	/*if(threadIdx.x == 0) {
		int len = tex1Dfetch(ellRowLenTex, idxX);
		for(int i=0;i<len;i++) {
			int2 X = tex2D(ellDataTex, idxX, i);
			dest[X.x] = __int_as_float(X.y);
		}
	}*/
	//__syncthreads();
}

#endif
